{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 0
   },
   "source": [
    "# Depthwise Convolution\n",
    ":label:`ch_depthwise_conv_cpu`\n",
    "\n",
    "In this section, we will follow the packing idea presented in :numref:`ch_packed_conv_cpu` to re-defined the computation of depthwise convolution and schedule it to run efficiently on CPUs. Similar to the 2-D convolution, tiling the data along the channel dimension and packing it into `NCHW[x]c` benefit the performance significantly.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 1,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "1"
    },
    "origin_pos": 1,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [],
   "source": [
    "import d2ltvm\n",
    "import numpy as np\n",
    "import tvm\n",
    "from tvm import te\n",
    "import timeit\n",
    "import os\n",
    "os.environ['KMP_AFFINITY']='granularity=fine,noduplicates,compact,1,0'\n",
    "\n",
    "target = 'llvm -mcpu=skylake-avx512'"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 2
   },
   "source": [
    "## Packing Data and Weight\n",
    "\n",
    "Recall the depthwise convolution described in :numref:`ch_depthwise_conv`, it differs from the 2-D convolution by having each channel of the input data convolved with a separated kernel. Therefore, the packing mechanism of input data is exactly the same as we did in :numref:`ch_packed_conv_cpu`. Kernel is a bit different, as the size is in `[oc, 1, kh, kw]`, which means that there is no need to tile the input channel.\n",
    "\n",
    "In other words, in the packing method below, we only pass one argument `c` to depict the channel, and another argument `tc` to depict the tiling size of channels. Other than that, it works very similarly as the `conv_pack` method defined in :numref:`ch_packed_conv_cpu`.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "2"
    },
    "origin_pos": 3,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [],
   "source": [
    "def depthwise_conv_pack(c, nh, nw, kh, kw, ph, pw, tc):\n",
    "    \"\"\"Pack data and weight for depthwise convolution\n",
    "       Note that the input channel of kernel is specified as 1,\n",
    "       and the output channel of kernel equals the input channel of data\n",
    "\n",
    "    c : input channel of data and output channel of kernel\n",
    "    nh, nw : input width and height\n",
    "    kh, kw : kernel width and height\n",
    "    ph, pw : height and width padding\n",
    "    tc : the tiling size of channels\n",
    "    \"\"\"\n",
    "    X = te.placeholder((c, nh, nw), name='X')\n",
    "    K = te.placeholder((c, 1, kh, kw), name='K')\n",
    "    PaddedX = d2ltvm.padding(X, ph, pw) if ph * pw != 0 else X\n",
    "    # make sure the channel tiling is valid\n",
    "    if c < tc:\n",
    "        tc = c\n",
    "    assert c % tc == 0\n",
    "    # pack X and K\n",
    "    PackedX = te.compute(\n",
    "        (c//tc, nh+ph*2, nw+pw*2, tc),\n",
    "        lambda c_out, x, y, c_in: PaddedX[c_out*tc + c_in, x, y],\n",
    "        name='PackedX')\n",
    "    PackedK = te.compute(\n",
    "        (c//tc, 1, kh, kw, 1, tc),\n",
    "        lambda c_out, _, x, y, __, c_in: K[\n",
    "            c_out*tc + c_in, 0, x, y],\n",
    "        name='PackedK')\n",
    "    return X, K, PaddedX, PackedX, PackedK"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 4
   },
   "source": [
    "## Computation\n",
    "\n",
    "Like in :numref:`ch_packed_conv_cpu`, we also need to re-implement the depthwise convolution computation accordingly.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 3,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "3"
    },
    "origin_pos": 5,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [],
   "source": [
    "def depthwise_conv(c, nh, nw, kh, kw, ph, pw, sh, sw, tc):\n",
    "    \"\"\"depthwise conv\n",
    "\n",
    "    c : number of channels for both input and output.\n",
    "    nh, nw : input width and height\n",
    "    kh, kw : kernel width and height\n",
    "    ph, pw : height and width padding\n",
    "    sh, sw : height and width strides\n",
    "    tc : the tiling sizes of channels\n",
    "    \"\"\"\n",
    "    X, K, PaddedX, PackedX, PackedK = depthwise_conv_pack(\n",
    "        c, nh, nw, kh, kw, ph, pw, tc)\n",
    "    # reduction axes\n",
    "    rkh = te.reduce_axis((0, kh), name='rkh')\n",
    "    rkw = te.reduce_axis((0, kw), name='rkw')\n",
    "    # output height and weights\n",
    "    oh = d2ltvm.conv_out_size(nh, kh, ph, sh)\n",
    "    ow = d2ltvm.conv_out_size(nw, kw, pw, sw)\n",
    "    # compute Y in the packed layout\n",
    "    PackedY = te.compute(\n",
    "        (c//tc, oh, ow, tc),\n",
    "        lambda c_out, x, y, c_in: te.sum(\n",
    "            (PackedX[c_out, x*sh+rkh, y*sw+rkw, c_in] *\n",
    "             PackedK[c_out, 0, rkh, rkw, 0, c_in]),\n",
    "            axis=[rkh, rkw]), name='PackedY')\n",
    "    \n",
    "    # Unpack the result\n",
    "    Y = te.compute((c, oh, ow),\n",
    "                    lambda c, x, y: PackedY[c//tc, x, y, c%tc],\n",
    "                    name='Y')\n",
    "    return X, K, Y, PaddedX, PackedX, PackedK, PackedY"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 6
   },
   "source": [
    "Let's quickly compile it using the default scheduling to compute the results.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 4,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "4"
    },
    "origin_pos": 7,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [],
   "source": [
    "c, n, k, p, s, tc = 32, 64, 3, 1, 1, 16\n",
    "\n",
    "X, K, Y, _, _, _, _ = depthwise_conv(c, n, n, k, k, p, p, s, s, tc)\n",
    "mod = tvm.build(te.create_schedule(Y.op), [X, K, Y])\n",
    "\n",
    "data, weight, out = d2ltvm.get_conv_data(c, c, n, k, p, s, tvm.nd.array, conv_type='depthwise')\n",
    "mod(data, weight, out)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 8
   },
   "source": [
    "And then verify the result.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 5,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "5"
    },
    "origin_pos": 9,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [],
   "source": [
    "data, weight, bias, out_mx = d2ltvm.get_conv_data_mxnet(c, c, n, k, p, s, conv_type='depthwise')\n",
    "d2ltvm.depthwise_conv_mxnet(data, weight, bias, out_mx, k, p, s)\n",
    "np.testing.assert_allclose(out_mx[0].asnumpy(), out.asnumpy(), atol=1e-5)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 10
   },
   "source": [
    "## Schedule\n",
    "\n",
    "The optimization strategy here is almost identical to `cache_block` defined in :numref:`ch_packed_conv_cpu`. The main  difference is in the channels, i.e. we don't need to reduce along the input channel dimension due to the compute nature of depthwise convolution.\n",
    "\n",
    "The tiling sizes for channel and width are set to make sure that the working set of the inner loop which calculates the cached output fits in the cache.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "6"
    },
    "origin_pos": 11,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [],
   "source": [
    "# tiling sizes for channel and width\n",
    "tc, tw = 16, 4\n",
    "\n",
    "def depthwise_cached_block(c, n, k, p, s):\n",
    "    X, K, Y, PaddedX, PackedX, PackedK, PackedY = depthwise_conv(\n",
    "        c, n, n, k, k, p, p, s, s, tc)\n",
    "    sch = te.create_schedule(Y.op)\n",
    "\n",
    "    CachedY = sch.cache_write(PackedY, 'global')\n",
    "\n",
    "    c_out, h, w, c_in = sch[PackedY].op.axis\n",
    "    w_out, w_in = sch[PackedY].split(w, factor=tw)\n",
    "    sch[PackedY].reorder(c_out, h, w_out, w_in, c_in)\n",
    "    c_out_h = sch[PackedY].fuse(c_out, h)\n",
    "    sch[PackedY].parallel(c_out_h)\n",
    "    sch[CachedY].compute_at(sch[PackedY], w_out)\n",
    "\n",
    "    cc_out, ch, cw, cc_in = sch[CachedY].op.axis\n",
    "    kh, kw = sch[CachedY].op.reduce_axis\n",
    "    sch[CachedY].reorder(cc_out, ch, kh, kw, cw, cc_in)\n",
    "    sch[CachedY].vectorize(cc_in)\n",
    "    sch[CachedY].unroll(cw)\n",
    "    \n",
    "    # Schedule the padding by adding thread-level parallelism\n",
    "    if PaddedX != X:\n",
    "        sch[PaddedX].parallel(PaddedX.op.axis[0])\n",
    "    # Optimize the packing of X and K\n",
    "    sch[PackedX].parallel(sch[PackedX].fuse(*PackedX.op.axis[0:2]))\n",
    "    sch[PackedX].unroll(PackedX.op.axis[-1])\n",
    "    sch[PackedK].parallel(sch[PackedK].fuse(*PackedK.op.axis[0:2]))\n",
    "    sch[PackedK].unroll(PackedK.op.axis[-1])\n",
    "    # Optimize the unpacking of Y\n",
    "    sch[Y].parallel(sch[Y].fuse(*Y.op.axis[0:2]))\n",
    "    sch[Y].unroll(Y.op.axis[-1])\n",
    "    return sch, (X, K, Y)\n",
    "\n",
    "# c, n, k, p, s were defined in the previous code block\n",
    "sch, args = depthwise_cached_block(c, n, k, p, s)\n",
    "# Uncomment the following line to see the long\n",
    "# psuedo codes because of unrolling.\n",
    "# tvm.lower(sch, args, simple_mode=True)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 12
   },
   "source": [
    "As the scheduling is vastly changed, let's do another round of sanity check.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 7,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "7"
    },
    "origin_pos": 13,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [],
   "source": [
    "mod = tvm.build(sch, args, target)\n",
    "ctx = tvm.context(target, 0)\n",
    "data, weight, out = d2ltvm.get_conv_data(\n",
    "            c, c, n, k, p, s, lambda x: tvm.nd.array(x, ctx=ctx), conv_type='depthwise')\n",
    "mod(data, weight, out)\n",
    "\n",
    "data, weight, bias, out_mx = d2ltvm.get_conv_data_mxnet(c, c, n, k, p, s, conv_type='depthwise')\n",
    "d2ltvm.depthwise_conv_mxnet(data, weight, bias, out_mx, k, p, s)\n",
    "np.testing.assert_allclose(out_mx[0].asnumpy(), out.asnumpy(), atol=1e-5)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 14
   },
   "source": [
    "## Benchmark\n",
    "\n",
    "Finally, let's benchmark the results against MXNet.\n",
    "\n",
    "The following benchmarking method is very similar to `bench_conv_tvm` defined in :numref:`ch_conv_cpu`, with two differences:\n",
    "\n",
    "1. The signature of the convolution functions (depthwise convolution only takes one channel input).\n",
    "2. The way to compute the FLOPs of computation (the input channel dimension of depthwise convolution is 1).\n",
    "\n",
    "We don't unify the benchmarking of depthwise convolution into the `bench_conv_tvm` method in order to reduce the possible confusion it may cause.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 8,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "8"
    },
    "origin_pos": 15,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [],
   "source": [
    "# Save to the d2ltvm package.\n",
    "def bench_depthwise_conv_tvm(func, sizes, target):\n",
    "    def workload(nrepeats):\n",
    "        timer = mod.time_evaluator(mod.entry_name, ctx=ctx, number=nrepeats)\n",
    "        return timer(x, k, y).mean * nrepeats\n",
    "    gflops, times = [], []\n",
    "    for (c, n, k) in sizes:\n",
    "        args = c, n, k, (k-1)//2, 1 # c, n, k, p, s\n",
    "        s, (X, K, Y) = func(*args)\n",
    "        mod = tvm.build(s, [X, K, Y], target)\n",
    "        ctx = tvm.context(target, 0)\n",
    "        x, k, y = d2ltvm.get_conv_data(\n",
    "            args[0], *args, lambda x: tvm.nd.array(x, ctx=ctx), conv_type='depthwise')\n",
    "        times.append(d2ltvm.bench_workload(workload))\n",
    "        gflops.append(d2ltvm.conv_gflop(1, *args))\n",
    "    return np.array(gflops) / np.array(times)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 16
   },
   "source": [
    "Similarly, the timing methods for depthwise convolution in MXNet are largely duplicated from the corresponding methods defined in :numref:`ch_conv_cpu`.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 9,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "9"
    },
    "origin_pos": 17,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [],
   "source": [
    "# Save to the d2ltvm package.\n",
    "def depthwise_conv_timer_mxnet(c, n, k, ctx):\n",
    "    \"\"\"Benchmark convolution in MXNet\n",
    "\n",
    "    c : input, output channels\n",
    "    n : input width and height\n",
    "    k : kernel width and height\n",
    "    \"\"\"\n",
    "    timer = timeit.Timer(\n",
    "        setup='import d2ltvm\\n'\n",
    "        'import mxnet as mx\\n'\n",
    "        'c, n, k, p, s = %d, %d, %d, %d, 1\\n'\n",
    "        'data, weight, bias, out = d2ltvm.get_conv_data_mxnet(\\n'\n",
    "        '    c, c, n, k, p, s, \"%s\", \"%s\")'%(c, n, k, (k-1)//2, ctx, 'depthwise'),\n",
    "        stmt='d2ltvm.depthwise_conv_mxnet(data, weight, bias, out, k, p, s);'\n",
    "        'out.wait_to_read()')\n",
    "    return timer.timeit\n",
    "\n",
    "# Save to the d2ltvm package.\n",
    "def bench_depthwise_conv_mxnet(sizes, ctx='cpu'):\n",
    "    \"\"\"Return the GFLOPS of MXNet convolution\"\"\"\n",
    "    return [d2ltvm.conv_gflop(1, c, n, k, (k-1)//2, 1) /\n",
    "            d2ltvm.bench_workload(depthwise_conv_timer_mxnet(c, n, k, ctx))\n",
    "            for c, n, k in sizes]"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 18
   },
   "source": [
    "Now, let's benchmark against our MXNet baseline. We see that our depthwise convolution performance consistently outperform MXNet. As depthwise convolution is a memory-bound operator, we see that the performance saturates after channel size gets to 128.\n"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 10,
   "metadata": {
    "attributes": {
     "classes": [],
     "id": "",
     "n": "10"
    },
    "origin_pos": 19,
    "tab": [
     "tvm"
    ]
   },
   "outputs": [
    {
     "data": {
      "image/svg+xml": [
       "<?xml version=\"1.0\" encoding=\"utf-8\" standalone=\"no\"?>\n",
       "<!DOCTYPE svg PUBLIC \"-//W3C//DTD SVG 1.1//EN\"\n",
       "  \"http://www.w3.org/Graphics/SVG/1.1/DTD/svg11.dtd\">\n",
       "<!-- Created with matplotlib (https://matplotlib.org/) -->\n",
       "<svg height=\"207.83625pt\" version=\"1.1\" viewBox=\"0 0 303.778125 207.83625\" width=\"303.778125pt\" xmlns=\"http://www.w3.org/2000/svg\" xmlns:xlink=\"http://www.w3.org/1999/xlink\">\n",
       " <metadata>\n",
       "  <rdf:RDF xmlns:cc=\"http://creativecommons.org/ns#\" xmlns:dc=\"http://purl.org/dc/elements/1.1/\" xmlns:rdf=\"http://www.w3.org/1999/02/22-rdf-syntax-ns#\">\n",
       "   <cc:Work>\n",
       "    <dc:type rdf:resource=\"http://purl.org/dc/dcmitype/StillImage\"/>\n",
       "    <dc:date>2020-10-13T10:38:51.863748</dc:date>\n",
       "    <dc:format>image/svg+xml</dc:format>\n",
       "    <dc:creator>\n",
       "     <cc:Agent>\n",
       "      <dc:title>Matplotlib v3.3.2, https://matplotlib.org/</dc:title>\n",
       "     </cc:Agent>\n",
       "    </dc:creator>\n",
       "   </cc:Work>\n",
       "  </rdf:RDF>\n",
       " </metadata>\n",
       " <defs>\n",
       "  <style type=\"text/css\">*{stroke-linecap:butt;stroke-linejoin:round;}</style>\n",
       " </defs>\n",
       " <g id=\"figure_1\">\n",
       "  <g id=\"patch_1\">\n",
       "   <path d=\"M 0 207.83625 \n",
       "L 303.778125 207.83625 \n",
       "L 303.778125 0 \n",
       "L 0 0 \n",
       "z\n",
       "\" style=\"fill:none;\"/>\n",
       "  </g>\n",
       "  <g id=\"axes_1\">\n",
       "   <g id=\"patch_2\">\n",
       "    <path d=\"M 45.478125 170.28 \n",
       "L 296.578125 170.28 \n",
       "L 296.578125 7.2 \n",
       "L 45.478125 7.2 \n",
       "z\n",
       "\" style=\"fill:#ffffff;\"/>\n",
       "   </g>\n",
       "   <g id=\"matplotlib.axis_1\">\n",
       "    <g id=\"xtick_1\">\n",
       "     <g id=\"line2d_1\">\n",
       "      <path clip-path=\"url(#pf93d416e5c)\" d=\"M 207.771827 170.28 \n",
       "L 207.771827 7.2 \n",
       "\" style=\"fill:none;stroke:#b0b0b0;stroke-linecap:square;stroke-width:0.8;\"/>\n",
       "     </g>\n",
       "     <g id=\"line2d_2\">\n",
       "      <defs>\n",
       "       <path d=\"M 0 0 \n",
       "L 0 3.5 \n",
       "\" id=\"me78072cfae\" style=\"stroke:#000000;stroke-width:0.8;\"/>\n",
       "      </defs>\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.8;\" x=\"207.771827\" xlink:href=\"#me78072cfae\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "     <g id=\"text_1\">\n",
       "      <!-- $\\mathdefault{10^{2}}$ -->\n",
       "      <g transform=\"translate(198.971827 184.878438)scale(0.1 -0.1)\">\n",
       "       <defs>\n",
       "        <path d=\"M 12.40625 8.296875 \n",
       "L 28.515625 8.296875 \n",
       "L 28.515625 63.921875 \n",
       "L 10.984375 60.40625 \n",
       "L 10.984375 69.390625 \n",
       "L 28.421875 72.90625 \n",
       "L 38.28125 72.90625 \n",
       "L 38.28125 8.296875 \n",
       "L 54.390625 8.296875 \n",
       "L 54.390625 0 \n",
       "L 12.40625 0 \n",
       "z\n",
       "\" id=\"DejaVuSans-49\"/>\n",
       "        <path d=\"M 31.78125 66.40625 \n",
       "Q 24.171875 66.40625 20.328125 58.90625 \n",
       "Q 16.5 51.421875 16.5 36.375 \n",
       "Q 16.5 21.390625 20.328125 13.890625 \n",
       "Q 24.171875 6.390625 31.78125 6.390625 \n",
       "Q 39.453125 6.390625 43.28125 13.890625 \n",
       "Q 47.125 21.390625 47.125 36.375 \n",
       "Q 47.125 51.421875 43.28125 58.90625 \n",
       "Q 39.453125 66.40625 31.78125 66.40625 \n",
       "z\n",
       "M 31.78125 74.21875 \n",
       "Q 44.046875 74.21875 50.515625 64.515625 \n",
       "Q 56.984375 54.828125 56.984375 36.375 \n",
       "Q 56.984375 17.96875 50.515625 8.265625 \n",
       "Q 44.046875 -1.421875 31.78125 -1.421875 \n",
       "Q 19.53125 -1.421875 13.0625 8.265625 \n",
       "Q 6.59375 17.96875 6.59375 36.375 \n",
       "Q 6.59375 54.828125 13.0625 64.515625 \n",
       "Q 19.53125 74.21875 31.78125 74.21875 \n",
       "z\n",
       "\" id=\"DejaVuSans-48\"/>\n",
       "        <path d=\"M 19.1875 8.296875 \n",
       "L 53.609375 8.296875 \n",
       "L 53.609375 0 \n",
       "L 7.328125 0 \n",
       "L 7.328125 8.296875 \n",
       "Q 12.9375 14.109375 22.625 23.890625 \n",
       "Q 32.328125 33.6875 34.8125 36.53125 \n",
       "Q 39.546875 41.84375 41.421875 45.53125 \n",
       "Q 43.3125 49.21875 43.3125 52.78125 \n",
       "Q 43.3125 58.59375 39.234375 62.25 \n",
       "Q 35.15625 65.921875 28.609375 65.921875 \n",
       "Q 23.96875 65.921875 18.8125 64.3125 \n",
       "Q 13.671875 62.703125 7.8125 59.421875 \n",
       "L 7.8125 69.390625 \n",
       "Q 13.765625 71.78125 18.9375 73 \n",
       "Q 24.125 74.21875 28.421875 74.21875 \n",
       "Q 39.75 74.21875 46.484375 68.546875 \n",
       "Q 53.21875 62.890625 53.21875 53.421875 \n",
       "Q 53.21875 48.921875 51.53125 44.890625 \n",
       "Q 49.859375 40.875 45.40625 35.40625 \n",
       "Q 44.1875 33.984375 37.640625 27.21875 \n",
       "Q 31.109375 20.453125 19.1875 8.296875 \n",
       "z\n",
       "\" id=\"DejaVuSans-50\"/>\n",
       "       </defs>\n",
       "       <use transform=\"translate(0 0.765625)\" xlink:href=\"#DejaVuSans-49\"/>\n",
       "       <use transform=\"translate(63.623047 0.765625)\" xlink:href=\"#DejaVuSans-48\"/>\n",
       "       <use transform=\"translate(128.203125 39.046875)scale(0.7)\" xlink:href=\"#DejaVuSans-50\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"xtick_2\">\n",
       "     <g id=\"line2d_3\">\n",
       "      <defs>\n",
       "       <path d=\"M 0 0 \n",
       "L 0 2 \n",
       "\" id=\"m78580ad476\" style=\"stroke:#000000;stroke-width:0.6;\"/>\n",
       "      </defs>\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"75.263612\" xlink:href=\"#m78580ad476\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"xtick_3\">\n",
       "     <g id=\"line2d_4\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"108.646359\" xlink:href=\"#m78580ad476\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"xtick_4\">\n",
       "     <g id=\"line2d_5\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"132.331794\" xlink:href=\"#m78580ad476\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"xtick_5\">\n",
       "     <g id=\"line2d_6\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"150.703645\" xlink:href=\"#m78580ad476\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"xtick_6\">\n",
       "     <g id=\"line2d_7\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"165.714541\" xlink:href=\"#m78580ad476\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"xtick_7\">\n",
       "     <g id=\"line2d_8\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"178.406072\" xlink:href=\"#m78580ad476\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"xtick_8\">\n",
       "     <g id=\"line2d_9\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"189.399976\" xlink:href=\"#m78580ad476\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"xtick_9\">\n",
       "     <g id=\"line2d_10\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"199.097287\" xlink:href=\"#m78580ad476\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"xtick_10\">\n",
       "     <g id=\"line2d_11\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"264.840009\" xlink:href=\"#m78580ad476\" y=\"170.28\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"text_2\">\n",
       "     <!-- Size -->\n",
       "     <g transform=\"translate(160.764062 198.556563)scale(0.1 -0.1)\">\n",
       "      <defs>\n",
       "       <path d=\"M 53.515625 70.515625 \n",
       "L 53.515625 60.890625 \n",
       "Q 47.90625 63.578125 42.921875 64.890625 \n",
       "Q 37.9375 66.21875 33.296875 66.21875 \n",
       "Q 25.25 66.21875 20.875 63.09375 \n",
       "Q 16.5 59.96875 16.5 54.203125 \n",
       "Q 16.5 49.359375 19.40625 46.890625 \n",
       "Q 22.3125 44.4375 30.421875 42.921875 \n",
       "L 36.375 41.703125 \n",
       "Q 47.40625 39.59375 52.65625 34.296875 \n",
       "Q 57.90625 29 57.90625 20.125 \n",
       "Q 57.90625 9.515625 50.796875 4.046875 \n",
       "Q 43.703125 -1.421875 29.984375 -1.421875 \n",
       "Q 24.8125 -1.421875 18.96875 -0.25 \n",
       "Q 13.140625 0.921875 6.890625 3.21875 \n",
       "L 6.890625 13.375 \n",
       "Q 12.890625 10.015625 18.65625 8.296875 \n",
       "Q 24.421875 6.59375 29.984375 6.59375 \n",
       "Q 38.421875 6.59375 43.015625 9.90625 \n",
       "Q 47.609375 13.234375 47.609375 19.390625 \n",
       "Q 47.609375 24.75 44.3125 27.78125 \n",
       "Q 41.015625 30.8125 33.5 32.328125 \n",
       "L 27.484375 33.5 \n",
       "Q 16.453125 35.6875 11.515625 40.375 \n",
       "Q 6.59375 45.0625 6.59375 53.421875 \n",
       "Q 6.59375 63.09375 13.40625 68.65625 \n",
       "Q 20.21875 74.21875 32.171875 74.21875 \n",
       "Q 37.3125 74.21875 42.625 73.28125 \n",
       "Q 47.953125 72.359375 53.515625 70.515625 \n",
       "z\n",
       "\" id=\"DejaVuSans-83\"/>\n",
       "       <path d=\"M 9.421875 54.6875 \n",
       "L 18.40625 54.6875 \n",
       "L 18.40625 0 \n",
       "L 9.421875 0 \n",
       "z\n",
       "M 9.421875 75.984375 \n",
       "L 18.40625 75.984375 \n",
       "L 18.40625 64.59375 \n",
       "L 9.421875 64.59375 \n",
       "z\n",
       "\" id=\"DejaVuSans-105\"/>\n",
       "       <path d=\"M 5.515625 54.6875 \n",
       "L 48.1875 54.6875 \n",
       "L 48.1875 46.484375 \n",
       "L 14.40625 7.171875 \n",
       "L 48.1875 7.171875 \n",
       "L 48.1875 0 \n",
       "L 4.296875 0 \n",
       "L 4.296875 8.203125 \n",
       "L 38.09375 47.515625 \n",
       "L 5.515625 47.515625 \n",
       "z\n",
       "\" id=\"DejaVuSans-122\"/>\n",
       "       <path d=\"M 56.203125 29.59375 \n",
       "L 56.203125 25.203125 \n",
       "L 14.890625 25.203125 \n",
       "Q 15.484375 15.921875 20.484375 11.0625 \n",
       "Q 25.484375 6.203125 34.421875 6.203125 \n",
       "Q 39.59375 6.203125 44.453125 7.46875 \n",
       "Q 49.3125 8.734375 54.109375 11.28125 \n",
       "L 54.109375 2.78125 \n",
       "Q 49.265625 0.734375 44.1875 -0.34375 \n",
       "Q 39.109375 -1.421875 33.890625 -1.421875 \n",
       "Q 20.796875 -1.421875 13.15625 6.1875 \n",
       "Q 5.515625 13.8125 5.515625 26.8125 \n",
       "Q 5.515625 40.234375 12.765625 48.109375 \n",
       "Q 20.015625 56 32.328125 56 \n",
       "Q 43.359375 56 49.78125 48.890625 \n",
       "Q 56.203125 41.796875 56.203125 29.59375 \n",
       "z\n",
       "M 47.21875 32.234375 \n",
       "Q 47.125 39.59375 43.09375 43.984375 \n",
       "Q 39.0625 48.390625 32.421875 48.390625 \n",
       "Q 24.90625 48.390625 20.390625 44.140625 \n",
       "Q 15.875 39.890625 15.1875 32.171875 \n",
       "z\n",
       "\" id=\"DejaVuSans-101\"/>\n",
       "      </defs>\n",
       "      <use xlink:href=\"#DejaVuSans-83\"/>\n",
       "      <use x=\"63.476562\" xlink:href=\"#DejaVuSans-105\"/>\n",
       "      <use x=\"91.259766\" xlink:href=\"#DejaVuSans-122\"/>\n",
       "      <use x=\"143.75\" xlink:href=\"#DejaVuSans-101\"/>\n",
       "     </g>\n",
       "    </g>\n",
       "   </g>\n",
       "   <g id=\"matplotlib.axis_2\">\n",
       "    <g id=\"ytick_1\">\n",
       "     <g id=\"line2d_12\">\n",
       "      <path clip-path=\"url(#pf93d416e5c)\" d=\"M 45.478125 84.584544 \n",
       "L 296.578125 84.584544 \n",
       "\" style=\"fill:none;stroke:#b0b0b0;stroke-linecap:square;stroke-width:0.8;\"/>\n",
       "     </g>\n",
       "     <g id=\"line2d_13\">\n",
       "      <defs>\n",
       "       <path d=\"M 0 0 \n",
       "L -3.5 0 \n",
       "\" id=\"mf16148eb24\" style=\"stroke:#000000;stroke-width:0.8;\"/>\n",
       "      </defs>\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.8;\" x=\"45.478125\" xlink:href=\"#mf16148eb24\" y=\"84.584544\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "     <g id=\"text_3\">\n",
       "      <!-- $\\mathdefault{10^{1}}$ -->\n",
       "      <g transform=\"translate(20.878125 88.383763)scale(0.1 -0.1)\">\n",
       "       <use transform=\"translate(0 0.684375)\" xlink:href=\"#DejaVuSans-49\"/>\n",
       "       <use transform=\"translate(63.623047 0.684375)\" xlink:href=\"#DejaVuSans-48\"/>\n",
       "       <use transform=\"translate(128.203125 38.965625)scale(0.7)\" xlink:href=\"#DejaVuSans-49\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_2\">\n",
       "     <g id=\"line2d_14\">\n",
       "      <defs>\n",
       "       <path d=\"M 0 0 \n",
       "L -2 0 \n",
       "\" id=\"m426e2736c2\" style=\"stroke:#000000;stroke-width:0.6;\"/>\n",
       "      </defs>\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"150.355616\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_3\">\n",
       "     <g id=\"line2d_15\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"134.640013\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_4\">\n",
       "     <g id=\"line2d_16\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"122.450044\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_5\">\n",
       "     <g id=\"line2d_17\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"112.490115\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_6\">\n",
       "     <g id=\"line2d_18\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"104.069115\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_7\">\n",
       "     <g id=\"line2d_19\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"96.774512\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_8\">\n",
       "     <g id=\"line2d_20\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"90.340217\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_9\">\n",
       "     <g id=\"line2d_21\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"46.719043\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_10\">\n",
       "     <g id=\"line2d_22\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"24.569145\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"ytick_11\">\n",
       "     <g id=\"line2d_23\">\n",
       "      <g>\n",
       "       <use style=\"stroke:#000000;stroke-width:0.6;\" x=\"45.478125\" xlink:href=\"#m426e2736c2\" y=\"8.853542\"/>\n",
       "      </g>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"text_4\">\n",
       "     <!-- GFLOPS -->\n",
       "     <g transform=\"translate(14.798438 108.219687)rotate(-90)scale(0.1 -0.1)\">\n",
       "      <defs>\n",
       "       <path d=\"M 59.515625 10.40625 \n",
       "L 59.515625 29.984375 \n",
       "L 43.40625 29.984375 \n",
       "L 43.40625 38.09375 \n",
       "L 69.28125 38.09375 \n",
       "L 69.28125 6.78125 \n",
       "Q 63.578125 2.734375 56.6875 0.65625 \n",
       "Q 49.8125 -1.421875 42 -1.421875 \n",
       "Q 24.90625 -1.421875 15.25 8.5625 \n",
       "Q 5.609375 18.5625 5.609375 36.375 \n",
       "Q 5.609375 54.25 15.25 64.234375 \n",
       "Q 24.90625 74.21875 42 74.21875 \n",
       "Q 49.125 74.21875 55.546875 72.453125 \n",
       "Q 61.96875 70.703125 67.390625 67.28125 \n",
       "L 67.390625 56.78125 \n",
       "Q 61.921875 61.421875 55.765625 63.765625 \n",
       "Q 49.609375 66.109375 42.828125 66.109375 \n",
       "Q 29.4375 66.109375 22.71875 58.640625 \n",
       "Q 16.015625 51.171875 16.015625 36.375 \n",
       "Q 16.015625 21.625 22.71875 14.15625 \n",
       "Q 29.4375 6.6875 42.828125 6.6875 \n",
       "Q 48.046875 6.6875 52.140625 7.59375 \n",
       "Q 56.25 8.5 59.515625 10.40625 \n",
       "z\n",
       "\" id=\"DejaVuSans-71\"/>\n",
       "       <path d=\"M 9.8125 72.90625 \n",
       "L 51.703125 72.90625 \n",
       "L 51.703125 64.59375 \n",
       "L 19.671875 64.59375 \n",
       "L 19.671875 43.109375 \n",
       "L 48.578125 43.109375 \n",
       "L 48.578125 34.8125 \n",
       "L 19.671875 34.8125 \n",
       "L 19.671875 0 \n",
       "L 9.8125 0 \n",
       "z\n",
       "\" id=\"DejaVuSans-70\"/>\n",
       "       <path d=\"M 9.8125 72.90625 \n",
       "L 19.671875 72.90625 \n",
       "L 19.671875 8.296875 \n",
       "L 55.171875 8.296875 \n",
       "L 55.171875 0 \n",
       "L 9.8125 0 \n",
       "z\n",
       "\" id=\"DejaVuSans-76\"/>\n",
       "       <path d=\"M 39.40625 66.21875 \n",
       "Q 28.65625 66.21875 22.328125 58.203125 \n",
       "Q 16.015625 50.203125 16.015625 36.375 \n",
       "Q 16.015625 22.609375 22.328125 14.59375 \n",
       "Q 28.65625 6.59375 39.40625 6.59375 \n",
       "Q 50.140625 6.59375 56.421875 14.59375 \n",
       "Q 62.703125 22.609375 62.703125 36.375 \n",
       "Q 62.703125 50.203125 56.421875 58.203125 \n",
       "Q 50.140625 66.21875 39.40625 66.21875 \n",
       "z\n",
       "M 39.40625 74.21875 \n",
       "Q 54.734375 74.21875 63.90625 63.9375 \n",
       "Q 73.09375 53.65625 73.09375 36.375 \n",
       "Q 73.09375 19.140625 63.90625 8.859375 \n",
       "Q 54.734375 -1.421875 39.40625 -1.421875 \n",
       "Q 24.03125 -1.421875 14.8125 8.828125 \n",
       "Q 5.609375 19.09375 5.609375 36.375 \n",
       "Q 5.609375 53.65625 14.8125 63.9375 \n",
       "Q 24.03125 74.21875 39.40625 74.21875 \n",
       "z\n",
       "\" id=\"DejaVuSans-79\"/>\n",
       "       <path d=\"M 19.671875 64.796875 \n",
       "L 19.671875 37.40625 \n",
       "L 32.078125 37.40625 \n",
       "Q 38.96875 37.40625 42.71875 40.96875 \n",
       "Q 46.484375 44.53125 46.484375 51.125 \n",
       "Q 46.484375 57.671875 42.71875 61.234375 \n",
       "Q 38.96875 64.796875 32.078125 64.796875 \n",
       "z\n",
       "M 9.8125 72.90625 \n",
       "L 32.078125 72.90625 \n",
       "Q 44.34375 72.90625 50.609375 67.359375 \n",
       "Q 56.890625 61.8125 56.890625 51.125 \n",
       "Q 56.890625 40.328125 50.609375 34.8125 \n",
       "Q 44.34375 29.296875 32.078125 29.296875 \n",
       "L 19.671875 29.296875 \n",
       "L 19.671875 0 \n",
       "L 9.8125 0 \n",
       "z\n",
       "\" id=\"DejaVuSans-80\"/>\n",
       "      </defs>\n",
       "      <use xlink:href=\"#DejaVuSans-71\"/>\n",
       "      <use x=\"77.490234\" xlink:href=\"#DejaVuSans-70\"/>\n",
       "      <use x=\"135.009766\" xlink:href=\"#DejaVuSans-76\"/>\n",
       "      <use x=\"187.097656\" xlink:href=\"#DejaVuSans-79\"/>\n",
       "      <use x=\"265.808594\" xlink:href=\"#DejaVuSans-80\"/>\n",
       "      <use x=\"326.111328\" xlink:href=\"#DejaVuSans-83\"/>\n",
       "     </g>\n",
       "    </g>\n",
       "   </g>\n",
       "   <g id=\"line2d_24\">\n",
       "    <path clip-path=\"url(#pf93d416e5c)\" d=\"M 56.891761 162.867273 \n",
       "L 113.959943 89.855853 \n",
       "L 171.028125 68.901284 \n",
       "L 228.096307 46.405338 \n",
       "L 285.164489 24.653258 \n",
       "\" style=\"fill:none;stroke:#1f77b4;stroke-dasharray:5.55,2.4;stroke-dashoffset:0;stroke-width:1.5;\"/>\n",
       "   </g>\n",
       "   <g id=\"line2d_25\">\n",
       "    <path clip-path=\"url(#pf93d416e5c)\" d=\"M 56.891761 44.786525 \n",
       "L 113.959943 29.792733 \n",
       "L 171.028125 16.797444 \n",
       "L 228.096307 15.729554 \n",
       "L 285.164489 14.612727 \n",
       "\" style=\"fill:none;stroke:#ff7f0e;stroke-linecap:square;stroke-width:1.5;\"/>\n",
       "   </g>\n",
       "   <g id=\"patch_3\">\n",
       "    <path d=\"M 45.478125 170.28 \n",
       "L 45.478125 7.2 \n",
       "\" style=\"fill:none;stroke:#000000;stroke-linecap:square;stroke-linejoin:miter;stroke-width:0.8;\"/>\n",
       "   </g>\n",
       "   <g id=\"patch_4\">\n",
       "    <path d=\"M 296.578125 170.28 \n",
       "L 296.578125 7.2 \n",
       "\" style=\"fill:none;stroke:#000000;stroke-linecap:square;stroke-linejoin:miter;stroke-width:0.8;\"/>\n",
       "   </g>\n",
       "   <g id=\"patch_5\">\n",
       "    <path d=\"M 45.478125 170.28 \n",
       "L 296.578125 170.28 \n",
       "\" style=\"fill:none;stroke:#000000;stroke-linecap:square;stroke-linejoin:miter;stroke-width:0.8;\"/>\n",
       "   </g>\n",
       "   <g id=\"patch_6\">\n",
       "    <path d=\"M 45.478125 7.2 \n",
       "L 296.578125 7.2 \n",
       "\" style=\"fill:none;stroke:#000000;stroke-linecap:square;stroke-linejoin:miter;stroke-width:0.8;\"/>\n",
       "   </g>\n",
       "   <g id=\"legend_1\">\n",
       "    <g id=\"patch_7\">\n",
       "     <path d=\"M 225.507812 165.28 \n",
       "L 289.578125 165.28 \n",
       "Q 291.578125 165.28 291.578125 163.28 \n",
       "L 291.578125 134.92375 \n",
       "Q 291.578125 132.92375 289.578125 132.92375 \n",
       "L 225.507812 132.92375 \n",
       "Q 223.507812 132.92375 223.507812 134.92375 \n",
       "L 223.507812 163.28 \n",
       "Q 223.507812 165.28 225.507812 165.28 \n",
       "z\n",
       "\" style=\"fill:#ffffff;opacity:0.8;stroke:#cccccc;stroke-linejoin:miter;\"/>\n",
       "    </g>\n",
       "    <g id=\"line2d_26\">\n",
       "     <path d=\"M 227.507812 141.022188 \n",
       "L 247.507812 141.022188 \n",
       "\" style=\"fill:none;stroke:#1f77b4;stroke-dasharray:5.55,2.4;stroke-dashoffset:0;stroke-width:1.5;\"/>\n",
       "    </g>\n",
       "    <g id=\"line2d_27\"/>\n",
       "    <g id=\"text_5\">\n",
       "     <!-- mxnet -->\n",
       "     <g transform=\"translate(255.507812 144.522188)scale(0.1 -0.1)\">\n",
       "      <defs>\n",
       "       <path d=\"M 52 44.1875 \n",
       "Q 55.375 50.25 60.0625 53.125 \n",
       "Q 64.75 56 71.09375 56 \n",
       "Q 79.640625 56 84.28125 50.015625 \n",
       "Q 88.921875 44.046875 88.921875 33.015625 \n",
       "L 88.921875 0 \n",
       "L 79.890625 0 \n",
       "L 79.890625 32.71875 \n",
       "Q 79.890625 40.578125 77.09375 44.375 \n",
       "Q 74.3125 48.1875 68.609375 48.1875 \n",
       "Q 61.625 48.1875 57.5625 43.546875 \n",
       "Q 53.515625 38.921875 53.515625 30.90625 \n",
       "L 53.515625 0 \n",
       "L 44.484375 0 \n",
       "L 44.484375 32.71875 \n",
       "Q 44.484375 40.625 41.703125 44.40625 \n",
       "Q 38.921875 48.1875 33.109375 48.1875 \n",
       "Q 26.21875 48.1875 22.15625 43.53125 \n",
       "Q 18.109375 38.875 18.109375 30.90625 \n",
       "L 18.109375 0 \n",
       "L 9.078125 0 \n",
       "L 9.078125 54.6875 \n",
       "L 18.109375 54.6875 \n",
       "L 18.109375 46.1875 \n",
       "Q 21.1875 51.21875 25.484375 53.609375 \n",
       "Q 29.78125 56 35.6875 56 \n",
       "Q 41.65625 56 45.828125 52.96875 \n",
       "Q 50 49.953125 52 44.1875 \n",
       "z\n",
       "\" id=\"DejaVuSans-109\"/>\n",
       "       <path d=\"M 54.890625 54.6875 \n",
       "L 35.109375 28.078125 \n",
       "L 55.90625 0 \n",
       "L 45.3125 0 \n",
       "L 29.390625 21.484375 \n",
       "L 13.484375 0 \n",
       "L 2.875 0 \n",
       "L 24.125 28.609375 \n",
       "L 4.6875 54.6875 \n",
       "L 15.28125 54.6875 \n",
       "L 29.78125 35.203125 \n",
       "L 44.28125 54.6875 \n",
       "z\n",
       "\" id=\"DejaVuSans-120\"/>\n",
       "       <path d=\"M 54.890625 33.015625 \n",
       "L 54.890625 0 \n",
       "L 45.90625 0 \n",
       "L 45.90625 32.71875 \n",
       "Q 45.90625 40.484375 42.875 44.328125 \n",
       "Q 39.84375 48.1875 33.796875 48.1875 \n",
       "Q 26.515625 48.1875 22.3125 43.546875 \n",
       "Q 18.109375 38.921875 18.109375 30.90625 \n",
       "L 18.109375 0 \n",
       "L 9.078125 0 \n",
       "L 9.078125 54.6875 \n",
       "L 18.109375 54.6875 \n",
       "L 18.109375 46.1875 \n",
       "Q 21.34375 51.125 25.703125 53.5625 \n",
       "Q 30.078125 56 35.796875 56 \n",
       "Q 45.21875 56 50.046875 50.171875 \n",
       "Q 54.890625 44.34375 54.890625 33.015625 \n",
       "z\n",
       "\" id=\"DejaVuSans-110\"/>\n",
       "       <path d=\"M 18.3125 70.21875 \n",
       "L 18.3125 54.6875 \n",
       "L 36.8125 54.6875 \n",
       "L 36.8125 47.703125 \n",
       "L 18.3125 47.703125 \n",
       "L 18.3125 18.015625 \n",
       "Q 18.3125 11.328125 20.140625 9.421875 \n",
       "Q 21.96875 7.515625 27.59375 7.515625 \n",
       "L 36.8125 7.515625 \n",
       "L 36.8125 0 \n",
       "L 27.59375 0 \n",
       "Q 17.1875 0 13.234375 3.875 \n",
       "Q 9.28125 7.765625 9.28125 18.015625 \n",
       "L 9.28125 47.703125 \n",
       "L 2.6875 47.703125 \n",
       "L 2.6875 54.6875 \n",
       "L 9.28125 54.6875 \n",
       "L 9.28125 70.21875 \n",
       "z\n",
       "\" id=\"DejaVuSans-116\"/>\n",
       "      </defs>\n",
       "      <use xlink:href=\"#DejaVuSans-109\"/>\n",
       "      <use x=\"97.412109\" xlink:href=\"#DejaVuSans-120\"/>\n",
       "      <use x=\"156.591797\" xlink:href=\"#DejaVuSans-110\"/>\n",
       "      <use x=\"219.970703\" xlink:href=\"#DejaVuSans-101\"/>\n",
       "      <use x=\"281.494141\" xlink:href=\"#DejaVuSans-116\"/>\n",
       "     </g>\n",
       "    </g>\n",
       "    <g id=\"line2d_28\">\n",
       "     <path d=\"M 227.507812 155.700312 \n",
       "L 247.507812 155.700312 \n",
       "\" style=\"fill:none;stroke:#ff7f0e;stroke-linecap:square;stroke-width:1.5;\"/>\n",
       "    </g>\n",
       "    <g id=\"line2d_29\"/>\n",
       "    <g id=\"text_6\">\n",
       "     <!-- tvm -->\n",
       "     <g transform=\"translate(255.507812 159.200312)scale(0.1 -0.1)\">\n",
       "      <defs>\n",
       "       <path d=\"M 2.984375 54.6875 \n",
       "L 12.5 54.6875 \n",
       "L 29.59375 8.796875 \n",
       "L 46.6875 54.6875 \n",
       "L 56.203125 54.6875 \n",
       "L 35.6875 0 \n",
       "L 23.484375 0 \n",
       "z\n",
       "\" id=\"DejaVuSans-118\"/>\n",
       "      </defs>\n",
       "      <use xlink:href=\"#DejaVuSans-116\"/>\n",
       "      <use x=\"39.208984\" xlink:href=\"#DejaVuSans-118\"/>\n",
       "      <use x=\"98.388672\" xlink:href=\"#DejaVuSans-109\"/>\n",
       "     </g>\n",
       "    </g>\n",
       "   </g>\n",
       "  </g>\n",
       " </g>\n",
       " <defs>\n",
       "  <clipPath id=\"pf93d416e5c\">\n",
       "   <rect height=\"163.08\" width=\"251.1\" x=\"45.478125\" y=\"7.2\"/>\n",
       "  </clipPath>\n",
       " </defs>\n",
       "</svg>\n"
      ],
      "text/plain": [
       "<Figure size 324x216 with 1 Axes>"
      ]
     },
     "metadata": {
      "needs_background": "light"
     },
     "output_type": "display_data"
    }
   ],
   "source": [
    "channels = 2**np.arange(4, 9)\n",
    "sizes = [(int(c), 64, 3) for c in channels] # a list of (c, n, k)\n",
    "tvm_gflops = bench_depthwise_conv_tvm(depthwise_cached_block, sizes, target)\n",
    "mxnet_gflops = bench_depthwise_conv_mxnet(sizes)\n",
    "d2ltvm.plot_gflops(channels, [mxnet_gflops, tvm_gflops], ['mxnet', 'tvm'])"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "origin_pos": 20
   },
   "source": [
    "## Summary\n",
    "\n",
    "- We can get good performance out of depthwise convolution by following the same rules of optimizing 2-D convolution.\n",
    "\n",
    "## Exercises\n",
    "\n",
    "- Try different tiling sizes.\n"
   ]
  }
 ],
 "metadata": {
  "language_info": {
   "name": "python"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 4
}